#include <cstdio>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>

#define THREAD_PER_BLOCK 256

__global__ void reduce_v0_shared(float *d_input, float *d_output)
{
    __shared__ float sdata[THREAD_PER_BLOCK];

    unsigned int tid = threadIdx.x;
    unsigned int gtid = blockDim.x * blockIdx.x + threadIdx.x;
    sdata[tid] = d_input[gtid];
    __syncthreads();

    for (unsigned int i = 1; i < blockDim.x; i *= 2)
    {
        if (tid % (i * 2) == 0)
        {
            sdata[tid] += sdata[tid + i];
        }
        __syncthreads();
    }

    if (tid == 0)
        d_output[blockIdx.x] = sdata[0];
}

bool check(float *result, float *output, int n)
{
    for (size_t i = 0; i < n; i++)
    {
        if (abs(result[i] - output[i]) >= 0.05)
            return false;
    }
    return true;
}

int main(int argc, char const *argv[])
{
    const int N = 32 * 1024 * 1024;
    int block_num = N / THREAD_PER_BLOCK;
    float *input = (float *)malloc(sizeof(float) * N);
    float *d_input;
    cudaMalloc((void **)&d_input, sizeof(float) * N);

    float *output = (float *)malloc(sizeof(float) * block_num);
    float *d_output;
    cudaMalloc((void **)&d_output, sizeof(float) * block_num);

    float *result = (float *)malloc(sizeof(float) * block_num);

    for (size_t i = 0; i < N; i++)
    {
        input[i] = 2.0 * (float)drand48() - 1.0;
    }
    for (size_t i = 0; i < block_num; i++)
    {
        float cur_res = 0;
        for (size_t j = 0; j < THREAD_PER_BLOCK; j++)
        {
            cur_res += input[i * THREAD_PER_BLOCK + j];
        }
        result[i] = cur_res;
    }

    cudaMemcpy(d_input, input, N * sizeof(float), cudaMemcpyHostToDevice);

    dim3 Grid(N / THREAD_PER_BLOCK, 1);
    dim3 Block(THREAD_PER_BLOCK, 1);

    reduce_v0_shared<<<Grid, Block>>>(d_input, d_output);

    cudaMemcpy(output, d_output, block_num * sizeof(float), cudaMemcpyDeviceToHost);

    if (check(result, output, block_num))
        printf("right\n");
    else
        printf("wrong\n");

    cudaFree(d_input);
    cudaFree(d_output);

    return 0;
}
